<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<title>Mali OpenCL SDK v1.1.0: From C to OpenCL</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="navtree.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="resize.js"></script>
<script type="text/javascript" src="navtree.js"></script>
<script type="text/javascript">
  $(document).ready(initResizable);
</script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
  $(document).ready(function() { searchBox.OnSelectItem(0); });
</script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
 <tbody>
 <tr style="height: 56px;">
  <td id="projectlogo"><img alt="Logo" src="logo.png"/></td>
  <td style="padding-left: 0.5em;">
   <div id="projectname">Mali OpenCL SDK v1.1.0
   </div>
  </td>
 </tr>
 </tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.2 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
  <div id="navrow1" class="tabs">
    <ul class="tablist">
      <li><a href="index.html"><span>Home</span></a></li>
      <li class="current"><a href="pages.html"><span>Help&#160;and&#160;Tutorials</span></a></li>
      <li><a href="files.html"><span>Files</span></a></li>
      <li>
        <div id="MSearchBox" class="MSearchBoxInactive">
        <span class="left">
          <img id="MSearchSelect" src="search/mag_sel.png"
               onmouseover="return searchBox.OnSearchSelectShow()"
               onmouseout="return searchBox.OnSearchSelectHide()"
               alt=""/>
          <input type="text" id="MSearchField" value="Search" accesskey="S"
               onfocus="searchBox.OnSearchFieldFocus(true)" 
               onblur="searchBox.OnSearchFieldFocus(false)" 
               onkeyup="searchBox.OnSearchFieldChange(event)"/>
          </span><span class="right">
            <a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
          </span>
        </div>
      </li>
    </ul>
  </div>
</div><!-- top -->
<div id="side-nav" class="ui-resizable side-nav-resizable">
  <div id="nav-tree">
    <div id="nav-tree-contents">
      <div id="nav-sync" class="sync"></div>
    </div>
  </div>
  <div id="splitbar" style="-moz-user-select:none;" 
       class="ui-resizable-handle">
  </div>
</div>
<script type="text/javascript">
$(document).ready(function(){initNavTree('hello_world_open_c_l_tutorial.html','');});
</script>
<div id="doc-content">
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
     onmouseover="return searchBox.OnSearchSelectShow()"
     onmouseout="return searchBox.OnSearchSelectHide()"
     onkeydown="return searchBox.OnSearchSelectKey(event)">
<a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(0)"><span class="SelectionMark">&#160;</span>All</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(1)"><span class="SelectionMark">&#160;</span>Classes</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(2)"><span class="SelectionMark">&#160;</span>Files</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(3)"><span class="SelectionMark">&#160;</span>Functions</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(4)"><span class="SelectionMark">&#160;</span>Variables</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(5)"><span class="SelectionMark">&#160;</span>Macros</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(6)"><span class="SelectionMark">&#160;</span>Pages</a></div>

<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0" 
        name="MSearchResults" id="MSearchResults">
</iframe>
</div>

<div class="header">
  <div class="headertitle">
<div class="title">From C to OpenCL </div>  </div>
</div><!--header-->
<div class="contents">
<div class="textblock"><p>How to move from a C/C++ 'for loop' to an OpenCL kernel.</p>
<h1><a class="anchor" id="helloWorldOpenCLAlgorithm"></a>
The Algorithm</h1>
<p>In this tutorial we are going to look at a very simple algorithm. Adding two arrays of numbers together element-by-element and storing the results in a third array:</p>
<p>C<sub>n</sub> = A<sub>n</sub> + B<sub>n</sub></p>
<p>Despite the simplistic nature of this algorithm, it can benefit from being implemented in OpenCL. Each element of the array can be calculated independently because there are no dependencies between elements in the array. This means that each element can easily be calculated in parallel. This kind of workload is ideal for OpenCL.</p>
<p>Unless otherwise noted, all code snippets below come from the OpenCL implementation in <a class="el" href="hello__world__opencl_8cpp.html">hello_world_opencl.cpp</a>.</p>
<h1><a class="anchor" id="helloWorldOpenCLOpenCLImplementation"></a>
C/C++ Implementation</h1>
<p>If we assume that we have three arrays (<a class="el" href="hello__world__c_8cpp.html">hello_world_c.cpp</a>):</p>
<div class="fragment"><div class="line">    <span class="comment">/* Number of elements in the arrays of input and output data. */</span></div>
<div class="line">    <span class="keywordtype">int</span> arraySize = 1000000;</div>
<div class="line"></div>
<div class="line">    <span class="comment">/* Arrays to hold the input and output data. */</span></div>
<div class="line">    <span class="keywordtype">int</span>* inputA = <span class="keyword">new</span> <span class="keywordtype">int</span>[arraySize];</div>
<div class="line">    <span class="keywordtype">int</span>* inputB = <span class="keyword">new</span> <span class="keywordtype">int</span>[arraySize];</div>
<div class="line">    <span class="keywordtype">int</span>* output = <span class="keyword">new</span> <span class="keywordtype">int</span>[arraySize];</div>
</div><!-- fragment --><p> Then an implementation in C/C++ is trivial (<a class="el" href="hello__world__c_8cpp.html">hello_world_c.cpp</a>): </p>
<div class="fragment"><div class="line">    <span class="keywordflow">for</span> (<span class="keywordtype">int</span> i = 0; i &lt; arraySize; i++)</div>
<div class="line">    {</div>
<div class="line">        output[i] = inputA[i] + inputB[i];</div>
<div class="line">    }</div>
</div><!-- fragment --><p> Discounting any compiler optimisations, this code will execute sequentially on a CPU. This means that each element of the array in calculated in series. An artificial dependency has been created between each calculation.</p>
<p>The runtime for this code is proportional to the size of the arrays.</p>
<h1><a class="anchor" id="helloWorldOpenCLImplementation"></a>
OpenCL Implementation</h1>
<h2><a class="anchor" id="helloWorldOpenCLBasics"></a>
The Basics</h2>
<ol type="1">
<li><p class="startli"><b>Move the parallelisable code into an OpenCL kernel</b></p>
<p class="startli">Take the parallel portion of the code (in our case, the for loop) and move it into and OpenCL kernel. For our code, with no optimisations, it looks like this (<a class="el" href="hello__world__opencl_8cl.html">hello_world_opencl.cl</a>):</p>
<div class="fragment"><div class="line">__kernel <span class="keywordtype">void</span> <a class="code" href="hello__world__opencl_8cl.html#a1e483e8dc82df49f47005c9ede1f0ed7" title="Hello World kernel function.">hello_world_opencl</a>(__global <span class="keywordtype">int</span>* restrict inputA,</div>
<div class="line">                                 __global <span class="keywordtype">int</span>* restrict inputB,</div>
<div class="line">                                 __global <span class="keywordtype">int</span>* restrict output)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * Set i to be the ID of the kernel instance.</span></div>
<div class="line"><span class="comment">     * If the global work size (set by clEnqueueNDRangeKernel) is n,</span></div>
<div class="line"><span class="comment">     * then n kernels will be run and i will be in the range [0, n - 1].</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    <span class="keywordtype">int</span> i = get_global_id(0);</div>
<div class="line"></div>
<div class="line">    <span class="comment">/* Use i as an index into the three arrays. */</span></div>
<div class="line">    output[i] = inputA[i] + inputB[i];</div>
<div class="line">}</div>
</div><!-- fragment --></li>
<li><p class="startli"><b>Run multiple instances of the kernel</b></p>
<p class="startli">There is no loop inside the kernel, so to make the code operate on all elements of the array, we must run several instances of the same kernel:</p>
<div class="fragment"><div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * Each instance of our OpenCL kernel operates on a single element of each array so the number of</span></div>
<div class="line"><span class="comment">     * instances needed is the number of elements in the array.</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    <span class="keywordtype">size_t</span> globalWorksize[1] = {arraySize};</div>
<div class="line">    <span class="comment">/* Enqueue the kernel */</span></div>
<div class="line">    <span class="keywordflow">if</span> (!<a class="code" href="common_8cpp.html#a73727ed6f5d821d8c97756404e145644" title="Check an OpenCL error number for errors.">checkSuccess</a>(clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, globalWorksize, NULL, 0, NULL, &amp;event)))</div>
<div class="line">    {</div>
<div class="line">        <a class="code" href="common_8cpp.html#acd459e74d0cef3c4616c7ce1a5a47f4d" title="Release any OpenCL objects that have been created.">cleanUpOpenCL</a>(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);</div>
<div class="line">        cerr &lt;&lt; <span class="stringliteral">&quot;Failed enqueuing the kernel. &quot;</span> &lt;&lt; __FILE__ &lt;&lt; <span class="stringliteral">&quot;:&quot;</span>&lt;&lt; __LINE__ &lt;&lt; endl;</div>
<div class="line">        <span class="keywordflow">return</span> 1;</div>
<div class="line">    }</div>
</div><!-- fragment --><p> This submits arraySize number of kernel instances to the OpenCL device.</p>
<p class="startli">Each instance is assigned a unique ID which we use to pick which element of the array each instance operates on (see the kernel above).</p>
<p class="startli">Because we've not specified any dependencies between kernels, the OpenCL device is free to run the instances of the kernel in parallel. The only limit on parallelism now is the device capabilities.</p>
<p class="startli">The runtime for this code is proportional to the size of the arrays divided by the number of kernel instances that can operate in parallel.</p>
</li>
</ol>
<h2><a class="anchor" id="helloWorldOpenCLSetup"></a>
OpenCL Setup</h2>
<p>There is some OpenCL setup required before the code above can be run. Take a look at <a class="el" href="hello__world__opencl_8cpp.html">hello_world_opencl.cpp</a> for more information.</p>
<h2><a class="anchor" id="helloWorldOpenCLMemory"></a>
Memory</h2>
<p>Because operations are now happening on the GPU rather than the CPU, we need to understand the location of any data we use. It is important to know whether the data is the GPU or CPU memory space.</p>
<p>In a desktop system the GPU and CPU have their own memories which are separated by a relatively slow bus. This can mean that sharing memory between the CPU and GPU is an expensive operation.</p>
<p>On most embedded systems with a Mali-T600 series GPU, the CPU and GPU share a common memory. It is therefore possible to share memory between the CPU and GPU relatively cheaply.</p>
<p>Because of these system differences, OpenCL supports many ways to allocate and share memory between devices.</p>
<p>Here is one way to share memory between devices which aims to eliminate copying memory from one device to another (in a shared memory system):</p>
<ol type="1">
<li><p class="startli"><b>Ask the OpenCL implementation to allocate some memory</b></p>
<p class="startli">In this sample we need three blocks of memory (two inputs and the output).</p>
<p class="startli">We use arrays in a C/C++ implementation. To allocate the arrays, we would do (<a class="el" href="hello__world__c_8cpp.html">hello_world_c.cpp</a>):</p>
<div class="fragment"><div class="line">    <span class="comment">/* Number of elements in the arrays of input and output data. */</span></div>
<div class="line">    <span class="keywordtype">int</span> arraySize = 1000000;</div>
<div class="line"></div>
<div class="line">    <span class="comment">/* Arrays to hold the input and output data. */</span></div>
<div class="line">    <span class="keywordtype">int</span>* inputA = <span class="keyword">new</span> <span class="keywordtype">int</span>[arraySize];</div>
<div class="line">    <span class="keywordtype">int</span>* inputB = <span class="keyword">new</span> <span class="keywordtype">int</span>[arraySize];</div>
<div class="line">    <span class="keywordtype">int</span>* output = <span class="keyword">new</span> <span class="keywordtype">int</span>[arraySize];</div>
</div><!-- fragment --><p> In OpenCL we use memory buffers which are just blocks of memory with a certain size. To allocate the buffers, we do:</p>
<div class="fragment"><div class="line">    <span class="comment">/* Number of elements in the arrays of input and output data. */</span></div>
<div class="line">    cl_int arraySize = 1000000;</div>
<div class="line"></div>
<div class="line">    <span class="comment">/* The buffers are the size of the arrays. */</span></div>
<div class="line">    <span class="keywordtype">size_t</span> bufferSize = arraySize * <span class="keyword">sizeof</span>(cl_int);</div>
<div class="line"></div>
<div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * Ask the OpenCL implementation to allocate buffers for the data.</span></div>
<div class="line"><span class="comment">     * We ask the OpenCL implemenation to allocate memory rather than allocating</span></div>
<div class="line"><span class="comment">     * it on the CPU to avoid having to copy the data later.</span></div>
<div class="line"><span class="comment">     * The read/write flags relate to accesses to the memory from within the kernel.</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    <span class="keywordtype">bool</span> createMemoryObjectsSuccess = <span class="keyword">true</span>;</div>
<div class="line"></div>
<div class="line">    memoryObjects[0] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, bufferSize, NULL, &amp;errorNumber);</div>
<div class="line">    createMemoryObjectsSuccess &amp;= <a class="code" href="common_8cpp.html#a73727ed6f5d821d8c97756404e145644" title="Check an OpenCL error number for errors.">checkSuccess</a>(errorNumber);</div>
<div class="line"></div>
<div class="line">    memoryObjects[1] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, bufferSize, NULL, &amp;errorNumber);</div>
<div class="line">    createMemoryObjectsSuccess &amp;= <a class="code" href="common_8cpp.html#a73727ed6f5d821d8c97756404e145644" title="Check an OpenCL error number for errors.">checkSuccess</a>(errorNumber);</div>
<div class="line"></div>
<div class="line">    memoryObjects[2] = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR, bufferSize, NULL, &amp;errorNumber);</div>
<div class="line">    createMemoryObjectsSuccess &amp;= <a class="code" href="common_8cpp.html#a73727ed6f5d821d8c97756404e145644" title="Check an OpenCL error number for errors.">checkSuccess</a>(errorNumber);</div>
<div class="line"></div>
<div class="line">    <span class="keywordflow">if</span> (!createMemoryObjectsSuccess)</div>
<div class="line">    {</div>
<div class="line">        <a class="code" href="common_8cpp.html#acd459e74d0cef3c4616c7ce1a5a47f4d" title="Release any OpenCL objects that have been created.">cleanUpOpenCL</a>(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);</div>
<div class="line">        cerr &lt;&lt; <span class="stringliteral">&quot;Failed to create OpenCL buffer. &quot;</span> &lt;&lt; __FILE__ &lt;&lt; <span class="stringliteral">&quot;:&quot;</span>&lt;&lt; __LINE__ &lt;&lt; endl;</div>
<div class="line">        <span class="keywordflow">return</span> 1;</div>
<div class="line">    }</div>
</div><!-- fragment --><p> Although this looks much more complex, there are only three OpenCL API calls. The difference is that here we are checking for errors (which is good practise) and in the C/C++ we are not.</p>
</li>
<li><p class="startli"><b>Map the memory to a local pointer</b></p>
<p class="startli">Now the memory has been allocated but only the OpenCL implementation knows where it is. To access the buffers on the CPU we map them to a pointer:</p>
<div class="fragment"><div class="line">    <span class="comment">/* Map the memory buffers created by the OpenCL implementation to pointers so we can access them on the CPU. */</span></div>
<div class="line">    <span class="keywordtype">bool</span> mapMemoryObjectsSuccess = <span class="keyword">true</span>;</div>
<div class="line"></div>
<div class="line">    cl_int* inputA = (cl_int*)clEnqueueMapBuffer(commandQueue, memoryObjects[0], CL_TRUE, CL_MAP_WRITE, 0, bufferSize, 0, NULL, NULL, &amp;errorNumber);</div>
<div class="line">    mapMemoryObjectsSuccess &amp;= <a class="code" href="common_8cpp.html#a73727ed6f5d821d8c97756404e145644" title="Check an OpenCL error number for errors.">checkSuccess</a>(errorNumber);</div>
<div class="line"></div>
<div class="line">    cl_int* inputB = (cl_int*)clEnqueueMapBuffer(commandQueue, memoryObjects[1], CL_TRUE, CL_MAP_WRITE, 0, bufferSize, 0, NULL, NULL, &amp;errorNumber);</div>
<div class="line">    mapMemoryObjectsSuccess &amp;= <a class="code" href="common_8cpp.html#a73727ed6f5d821d8c97756404e145644" title="Check an OpenCL error number for errors.">checkSuccess</a>(errorNumber);</div>
<div class="line"></div>
<div class="line">    <span class="keywordflow">if</span> (!mapMemoryObjectsSuccess)</div>
<div class="line">    {</div>
<div class="line">       <a class="code" href="common_8cpp.html#acd459e74d0cef3c4616c7ce1a5a47f4d" title="Release any OpenCL objects that have been created.">cleanUpOpenCL</a>(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);</div>
<div class="line">       cerr &lt;&lt; <span class="stringliteral">&quot;Failed to map buffer. &quot;</span> &lt;&lt; __FILE__ &lt;&lt; <span class="stringliteral">&quot;:&quot;</span>&lt;&lt; __LINE__ &lt;&lt; endl;</div>
<div class="line">       <span class="keywordflow">return</span> 1;</div>
<div class="line">    }</div>
</div><!-- fragment --><p> These pointers can now be used as normal C/C++ pointers.</p>
</li>
<li><p class="startli"><b>Initialise the data on the CPU</b></p>
<p class="startli">Because we have pointers to the memory, this step is the same as on the CPU:</p>
<div class="fragment"><div class="line">    <span class="keywordflow">for</span> (<span class="keywordtype">int</span> i = 0; i &lt; arraySize; i++)</div>
<div class="line">    {</div>
<div class="line">       inputA[i] = i;</div>
<div class="line">       inputB[i] = i;</div>
<div class="line">    }</div>
</div><!-- fragment --></li>
<li><p class="startli"><b>Un-map the buffers</b></p>
<p class="startli">To allow the OpenCL device to use the buffers we must un-map them from the CPU:</p>
<div class="fragment"><div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * Unmap the memory objects as we have finished using them from the CPU side.</span></div>
<div class="line"><span class="comment">     * We unmap the memory because otherwise:</span></div>
<div class="line"><span class="comment">     * - reads and writes to that memory from inside a kernel on the OpenCL side are undefined.</span></div>
<div class="line"><span class="comment">     * - the OpenCL implementation cannot free the memory when it is finished.</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    <span class="keywordflow">if</span> (!<a class="code" href="common_8cpp.html#a73727ed6f5d821d8c97756404e145644" title="Check an OpenCL error number for errors.">checkSuccess</a>(clEnqueueUnmapMemObject(commandQueue, memoryObjects[0], inputA, 0, NULL, NULL)))</div>
<div class="line">    {</div>
<div class="line">       <a class="code" href="common_8cpp.html#acd459e74d0cef3c4616c7ce1a5a47f4d" title="Release any OpenCL objects that have been created.">cleanUpOpenCL</a>(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);</div>
<div class="line">       cerr &lt;&lt; <span class="stringliteral">&quot;Unmapping memory objects failed &quot;</span> &lt;&lt; __FILE__ &lt;&lt; <span class="stringliteral">&quot;:&quot;</span>&lt;&lt; __LINE__ &lt;&lt; endl;</div>
<div class="line">       <span class="keywordflow">return</span> 1;</div>
<div class="line">    }</div>
<div class="line"></div>
<div class="line">    <span class="keywordflow">if</span> (!<a class="code" href="common_8cpp.html#a73727ed6f5d821d8c97756404e145644" title="Check an OpenCL error number for errors.">checkSuccess</a>(clEnqueueUnmapMemObject(commandQueue, memoryObjects[1], inputB, 0, NULL, NULL)))</div>
<div class="line">    {</div>
<div class="line">       <a class="code" href="common_8cpp.html#acd459e74d0cef3c4616c7ce1a5a47f4d" title="Release any OpenCL objects that have been created.">cleanUpOpenCL</a>(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);</div>
<div class="line">       cerr &lt;&lt; <span class="stringliteral">&quot;Unmapping memory objects failed &quot;</span> &lt;&lt; __FILE__ &lt;&lt; <span class="stringliteral">&quot;:&quot;</span>&lt;&lt; __LINE__ &lt;&lt; endl;</div>
<div class="line">       <span class="keywordflow">return</span> 1;</div>
<div class="line">    }</div>
</div><!-- fragment --></li>
<li><p class="startli"><b>Map the data to the kernels</b></p>
<p class="startli">We have to tell the kernels which data to use for its inputs before we schedule it to run.</p>
<p class="startli">Here, we map the memory objects to the parameters of the OpenCL kernel:</p>
<div class="fragment"><div class="line">    <span class="keywordtype">bool</span> setKernelArgumentsSuccess = <span class="keyword">true</span>;</div>
<div class="line">    setKernelArgumentsSuccess &amp;= <a class="code" href="common_8cpp.html#a73727ed6f5d821d8c97756404e145644" title="Check an OpenCL error number for errors.">checkSuccess</a>(clSetKernelArg(kernel, 0, <span class="keyword">sizeof</span>(cl_mem), &amp;memoryObjects[0]));</div>
<div class="line">    setKernelArgumentsSuccess &amp;= <a class="code" href="common_8cpp.html#a73727ed6f5d821d8c97756404e145644" title="Check an OpenCL error number for errors.">checkSuccess</a>(clSetKernelArg(kernel, 1, <span class="keyword">sizeof</span>(cl_mem), &amp;memoryObjects[1]));</div>
<div class="line">    setKernelArgumentsSuccess &amp;= <a class="code" href="common_8cpp.html#a73727ed6f5d821d8c97756404e145644" title="Check an OpenCL error number for errors.">checkSuccess</a>(clSetKernelArg(kernel, 2, <span class="keyword">sizeof</span>(cl_mem), &amp;memoryObjects[2]));</div>
<div class="line"></div>
<div class="line">    <span class="keywordflow">if</span> (!setKernelArgumentsSuccess)</div>
<div class="line">    {</div>
<div class="line">        <a class="code" href="common_8cpp.html#acd459e74d0cef3c4616c7ce1a5a47f4d" title="Release any OpenCL objects that have been created.">cleanUpOpenCL</a>(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);</div>
<div class="line">        cerr &lt;&lt; <span class="stringliteral">&quot;Failed setting OpenCL kernel arguments. &quot;</span> &lt;&lt; __FILE__ &lt;&lt; <span class="stringliteral">&quot;:&quot;</span>&lt;&lt; __LINE__ &lt;&lt; endl;</div>
<div class="line">        <span class="keywordflow">return</span> 1;</div>
<div class="line">    }</div>
</div><!-- fragment --></li>
<li><p class="startli"><b>Run the kernels</b></p>
<p class="startli">For the kernel code and how to schedule it, see <a class="el" href="hello_world_open_c_l_tutorial.html#helloWorldOpenCLBasics">The Basics</a>.</p>
</li>
<li><p class="startli"><b>Get the results</b></p>
<p class="startli">Once the calculations are complete we map the output buffer in the same way we mapped the input buffers. We can then read the results using the pointer as normal and then unmap the buffer as before.</p>
</li>
</ol>
<h1><a class="anchor" id="helloWorldOpenCLMoreInformation"></a>
More Information</h1>
<p>For more information have a look at the C/C++ code in <a class="el" href="hello__world__c_8cpp.html">hello_world_c.cpp</a> and the OpenCL version in <a class="el" href="hello__world__opencl_8cpp.html">hello_world_opencl.cpp</a> and <a class="el" href="hello__world__opencl_8cl.html">hello_world_opencl.cl</a>.</p>
<p>Main Tutorial: <a class="el" href="hello_world_tutorial.html">Hello World</a>.</p>
<p>Next section: <a class="el" href="hello_world_vector_tutorial.html">Vectorizing your OpenCL code</a>.</p>
<p><a class="el" href="hello_world_running_tutorial.html">Running the Samples</a> </p>
</div></div><!-- contents -->
</div><!-- doc-content -->
<!-- start footer part -->
<div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
  <ul>
    <li class="navelem"><a class="el" href="tutorials.html">Tutorials</a></li><li class="navelem"><a class="el" href="hello_world_tutorial.html">Hello World</a></li>
    <li class="footer">
        <a href="http://www.arm.com/">(C) ARM Ltd. 2013</a>
    </li>
  </ul>
</div>
</body>
</html>
